/* Begin template: kernel_call , file: kernel_call.tmpl */
<%!
       from Backends.Common.TemplateEngine.Functions import write_as_function_call, is_array, is_pointer, write_as_parameter_definition
%>
<%namespace name="declarations" file="device_declaration.tmpl"/>
<%namespace name="transfers" file="device_transfers.tmpl"/>


#define LLC_REDUCTION_FUNC(dest, source) dest = dest + source
#define CUDA_NUM_THREADS 64

/* Kernel initialization and call */
void call_${funcName}(
                  ${', '.join(str(var.declaration) + str(var) for var in parameters)}
) {
         int dimA = * dimA;
         int numThreadsPerBlock = CUDA_NUM_THREADS;
         int numBlocks = dimA / numThreadsPerBlock + (dimA % numThreadsPerBlock?1:0);
         int numElems = numBlocks * numThreadsPerBlock;
         int memSize = numElems * sizeof(double);
         dim3 dimGrid, dimBlock;

         /* Variable declaration */
         % for var in reduction_vars:
                   ${var.type} * reduction_loc_${var.name};
                   ${var.type} * reduction_cu_${var.name};
         % endfor

         /* Complex types need to be created on device memory */
         ${declarations.write_declarations(device_vars, '', '_cu')}     

         /* Initialization of reduction array */
         % for var in reduction_vars:
          reduction_loc_${var.name} = (${var.type} *) malloc( numElems * sizeof(${var.type}));
          /* TODO: Error check for valid mallocs */
          cudaMalloc((void **) (&reduction_cu_${var.name}), numElems * sizeof(${var.type}));
          /* This may be incorrect in case reduction don't start with 0 or 1 */
          // Incorrect: cudaMemset(reduction_cu_${var.name}, 0, numElems * sizeof(${var.type}));
         % endfor

                
         /* Initialization of declarations */
         ${declarations.write_allocate_to_device(device_vars, '', '_cu')}     
         ${transfers.write_transfer_to_device(device_vars, '', '_cu')}     

         /* Kernel Launch */
         dimGrid.x = numBlocks;
         dimBlock.x = numThreadsPerBlock;

         ${kernelName} <<< dimGrid, dimBlock >>> (${', '.join("reduction_cu_" + var.name for var in reduction_vars)}
         %if len(reduction_vars) > 0 and len(shared_vars) > 0:
               ,
         %endif 
         ${write_as_function_call(shared_vars, value = True)});

        checkCUDAError("memcpy");

        <%namespace name="snippets" file="snippets.tmpl"/>
        ${snippets.write_retrieve_all_data_from_device(reduction_vars, device_vars)}     

        % for var in reduction_vars:
            (*${var}) = kernelReduction_${var.type}(reduction_cu_${var.name}, numElems, (*${var.name}));
        % endfor

        /* By default, omp for has a wait at the end */
        % if not nowait:
            cudaThreadSynchronize();
        % endif

        % for var in reduction_vars:
        	cudaFree(reduction_cu_${var.name});
        % endfor

   }


/* End template: kernel_call , file: kernel_call.tmpl */
## ex: set filetype=c
